/** File name: cellRiver2.cu
*** Author: Tran Van Hoang
*** Date: 25/5/2015
*** Description: This version is developed on the previous version (cellRiver.cu). 
***              However, the pileline technique is used here.
****/

#include <sys/unistd.h>
#include <pthread.h>
#include <stdlib.h>
#include <stdio.h>
#include <cellNetwork.h>
#include <view.h>
#include "data.cu"
#include "network.cu"
#include <Xlib.h>

#define ROUNDS 20
#define THREADS 3
#define T1	0
#define T2	1
#define T3	2
#define BLSIZE 512
#define N 601

///Displaying Window
#define WINDOW_TITLE "CellRiver Pipeline version"
int plotSize, NSize;
Colormap colormap;
Display *dsp;	
int screenNumber;
unsigned long white;
unsigned long black;	
Window win;
GC gc;	
long eventMask;
XEvent evt;

//Main declaration
void *threadStep(void * args);

typedef struct {
	int threadId;
	NodeState *dataIn, *dataOut;
	canaux *channelIn, *channelOut;
	}data_thread;

data_thread  data[THREADS];

extern __global__ void compute(NodeState *tabi,NodeState *tabout, canaux *channels);

extern __global__ void SetZero(NodeState *tabi, int node_number);

void displayState(NodeState *nodeStates)
{
	int density;

    	for (int j = 0; j < NODE_NUMBER; j++)
    	{
	    density = (int) (nodeStates[j].density * 100.0);
	    if (density > 0.0)
	    {
    	    	showCellDensity(Cells[j], dsp, colormap, win, gc, plotSize, density);
	    }
    	}
}

void dataMove()
{
     	NodeState *tmpState;
     	canaux *tmpChannel;

        // Exchanging data between threads
	tmpState = data[0].dataOut;
	data[0].dataOut = data[1].dataIn;
	data[1].dataIn = tmpState;

	tmpState = data[1].dataOut;
	data[1].dataOut = data[2].dataIn;
	data[2].dataIn = tmpState;

	tmpChannel = data[0].channelOut;
	data[0].channelOut = data[1].channelIn;
	data[1].channelIn = tmpChannel;
}

void init_data_thread(data_thread *dt, int id)
{
	int i,err,nblocs;
        err= 0;
        nblocs = NODE_NUMBER / BLSIZE + 1;
	dt->threadId = id;

	switch (id)
		{
		case T1 : 
			dt->dataIn = (NodeState*) malloc(sizeof (NodeState) * NODE_NUMBER);
			//dt->channels = (canaux*) malloc(sizeof (canaux) * NODE_NUMBER);
			dt->channelIn = channels_h;//channels_h comes from network.cu
			err = cudaMalloc(&dt->dataOut,  sizeof (NodeState) * NODE_NUMBER);
			err = cudaMalloc(&dt->channelOut,  sizeof (canaux) * NODE_NUMBER);
			initState(dt->dataIn, NODE_NUMBER);
			break;
		case T2 : 
			err = cudaMalloc(&dt->dataIn,  sizeof (NodeState) * NODE_NUMBER);
			err |= cudaMalloc(&dt->dataOut,  sizeof (NodeState) * NODE_NUMBER);
			err = cudaMalloc(&dt->channelIn,  sizeof (canaux) * NODE_NUMBER);
  			SetZero<<<nblocs,BLSIZE>>>(dt->dataIn, NODE_NUMBER);
  			SetZero<<<nblocs,BLSIZE>>>(dt->dataOut, NODE_NUMBER);
			break;
		case T3 : 
			cudaMalloc(&dt->dataIn,  sizeof (NodeState) * NODE_NUMBER);
  			SetZero<<<nblocs,BLSIZE>>>(dt->dataIn, NODE_NUMBER);
			dt->dataOut = (NodeState *) malloc(sizeof (NodeState) * NODE_NUMBER);
			for (i=0; i < NODE_NUMBER; i++) 
			{
			    //dt->dataOut[i] = 0 ;
			}
			break;
		}
}

void free_data(data_thread dt)
{
        int id;
	id = dt.threadId; 

	switch (id)
		{
		case T1 : 
			free(dt.dataIn);
			cudaFree(dt.dataOut);
			//free(dt.channelIn);
			cudaFree(dt.channelOut);
			break;
		case T2 :
			cudaFree(dt.dataIn);
			cudaFree(dt.dataOut);
			cudaFree(dt.channelIn);
			break;
		case T3 : 
			cudaFree(dt.dataIn);
			free(dt.dataOut);
			break;
		}
}

void initWindow()
{
	plotSize = 1;
	NSize = 256;
	dsp = XOpenDisplay(NULL);	
	if (!dsp){ return;}
    	colormap = DefaultColormap(dsp, 0);

	screenNumber = DefaultScreen(dsp);
	white = WhitePixel(dsp, screenNumber);
	black = BlackPixel(dsp, screenNumber);	

	win = XCreateSimpleWindow(dsp,
			DefaultRootWindow(dsp),
			50, 50,// origin
			NSize*plotSize*3, NSize*plotSize*3,// size
			0, black,//border
			white//background
	);

	XStoreName(dsp, win, WINDOW_TITLE);
	XMapWindow(dsp, win);
	
	gc = XCreateGC(dsp, win, 0, NULL);	
	
	eventMask = StructureNotifyMask;
	XSelectInput(dsp, win, eventMask);
	do
	{
		XNextEvent(dsp, &evt);
	}while(evt.type != MapNotify);	
}

void event()
{
	eventMask = ButtonPressMask | ButtonReleaseMask;

	XSelectInput(dsp, win, eventMask);
	do{
		XNextEvent(dsp, &evt);
	}while(evt.type != ButtonRelease);
}

int main(int argc, char * argv[])
{
    	int i;
    	int steps, maxsteps;
    	void *status;

    	if (argc > 1)
	   maxsteps = atoi(argv[1]);
    	else
           maxsteps = ROUNDS;

    	pthread_t *threads;
    	pthread_attr_t attr;

    	threads=(pthread_t*) malloc( sizeof(pthread_t) * THREADS);
    	pthread_attr_init(&attr);
    	pthread_attr_setdetachstate(&attr, PTHREAD_CREATE_JOINABLE);

    	for (i=0; i<THREADS;i++)
    	{
    	    init_data_thread(&data[i], i);
    	}

    	//Initilizating window
    	initWindow();
    	int density;

	for (int j = 0; j < NODE_NUMBER; j++)
	{
	     density = (int) (data[0].dataIn[j].density * 100.0);
	     showCellDensity(Cells[j], dsp, colormap, win, gc, plotSize, density);
	}

	printf("\n-- Click on the window to continue!\n");
	event();
    	//Simulation
    	for (steps = 0; steps < maxsteps; steps++)
    	{
            for (i=0; i<THREADS; i++)
       	    {
           	pthread_create(&threads[i], NULL, threadStep, (void*) &(data[i]));
            }
            for (i = 0; i < THREADS; i++)
            {
           	pthread_join(threads[i], &status);
           	if(status);
               //Error;
            }

            dataMove();
	    //preparing for next step
	    NodeState *tmpState;
	    if (steps >= 2)
	    {
	    	tmpState = data[1].dataIn;
	    	data[1].dataIn = data[1].dataOut;
	    	data[1].dataOut = tmpState;
	    }
    	}

    	for (i=0; i < THREADS; i++) 
    	{
	    free_data (data[i]);
    	}

    	cudaDeviceReset();
	
	printf("\n-- Click on window to exit!\n");

	event();

    	printf("Pipe threads %d had finished for size N=%d ROUNDS=%d\n",THREADS,N, maxsteps);

    	return 0;
}

__global__ void SetZero(NodeState *nodeStates, int node_number)
{
    	int i = blockDim.x * blockIdx.x + threadIdx.x;

    	if (i < node_number)
    	{
	   nodeStates[i].density = (float) 0.0;
    	}
}


void toGPU(NodeState *tin, NodeState *tout, canaux *channels_h, canaux *channels_d)
{
	cudaMemcpy(tout, tin, NODE_NUMBER * sizeof(NodeState), cudaMemcpyHostToDevice);
	cudaMemcpy(channels_d, channels_h, NODE_NUMBER * sizeof(canaux), cudaMemcpyHostToDevice);
}

void processOnGPU(NodeState *tin, NodeState *tout, canaux *channels)
{
	int nblocs;
	nblocs = NODE_NUMBER / BLSIZE + 1;
	compute<<<nblocs,BLSIZE>>>(tin, tout, channels);
}

void fromGPU(NodeState *tin, NodeState *tout)
{
	cudaMemcpy(tout, tin, NODE_NUMBER * sizeof(NodeState), cudaMemcpyDeviceToHost);
        displayState(tout);
}

void *threadStep(void *args)
{
    	int id;
    	data_thread *sdata;
    	sdata = (data_thread*) args;
    	id  =  sdata->threadId;

    	switch (id) {
	      case T1: toGPU(sdata->dataIn, sdata->dataOut, sdata->channelIn, sdata->channelOut); break;
	      case T2: processOnGPU(sdata->dataIn, sdata->dataOut, sdata->channelIn); break;
	      case T3: fromGPU(sdata->dataIn, sdata->dataOut); break;
	}

    	return 0;
}

__global__ void compute(NodeState *nowStates, NodeState *nextStates, canaux *channels)
{
    	int idx = blockDim.x * blockIdx.x + threadIdx.x;

    	if (idx < NODE_NUMBER) 
    	{
	   //checking the neighbours

	    int nbIn = channels[idx].nbIn;
	    int nodeIn;
	    float receive = (float)0.0;

	    for (int i = 0; i < nbIn; i++)
	    {
	    	nodeIn = channels[idx].read[i].node;
	    	receive = receive + ((nowStates[nodeIn].density / 2.0) / (float) channels[nodeIn].nbIn);
	    }

	    nextStates[idx].density = (nowStates[idx].density / 2.0) + receive;
    	}
}

